__kernel void scan(__global int * g_in,
                   __global int * g_out,
                   __local int *l_temp,
                   int n) {

  int thid = get_local_id(0);
  int pout = 0, pin = 1;
  
  // Load input into shared memory.
  // This is exclusive scan, so shift right by one
  // and set first element to 0
  l_temp[pout*n + thid] = (thid > 0) ? g_in[thid-1] : 0;
  
  barrier(CLK_LOCAL_MEM_FENCE);
  for (int offset = 1; offset < n; offset *= 2)
  {
    pout = 1 - pout; // swap double buffer indices
    pin = 1 - pout;
    
    if (thid >= offset)
      l_temp[pout*n+thid] = l_temp[pin*n+thid] + l_temp[pin*n+thid - offset];
    else
      l_temp[pout*n+thid] = l_temp[pin*n+thid];
    
    barrier(CLK_LOCAL_MEM_FENCE);
  }
  
  // Copy output data
  g_out [thid] = l_temp[pout*n + thid];
}